/*
 * ocl.cpp
 *
 *  Created on: 15 cze 2014
 *      Author: Paweł Macioł
 */
#include "dbg.h"
#include "fv_threads.h"
#include "fv_helper_routines.h"
#include "fv_dictstr.h"
#include "ocl.h"
#include "Mesh.h"
#include "Field.h"
#include <memory>
#include <vector>
#include <limits>
#include <string>
#include <fstream>
#include <iostream>
// OpenCL C++ wraper
#define  __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>
#ifdef UNIX
       #include <GL/glx.h>
#endif


/* Paths to sources of kernels */
const char* kernels[] = {
		"/home/dwg/development/mod_fem_code/src/mod_fem_viewer/FemViewer/ocl/kernel_zero.cl"
		"/home/dwg/development/mod_fem_code/src/mod_fem_viewer/FemViewer/ocl/kernels.cl"
};

/* Registered types od compute devices */
static const
cl_device_type _devices[] = {
		CL_DEVICE_TYPE_CPU,
		CL_DEVICE_TYPE_GPU,
		CL_DEVICE_TYPE_ACCELERATOR
};


class ocl
{
public:
	enum eErr {
		TEST_FAILED = -100,
	};
	typedef cl::Platform 		oclPlatform;
	typedef cl::Device   		oclDevice;
	typedef cl::Context  		oclContext;
	typedef cl::CommandQueue	oclCmdQ;
	typedef cl::Program		 	oclProgram;
	typedef cl::Kernel	 		oclKernel;
	typedef cl::Event		 	oclEvent;
	typedef cl::Error			oclError;
	static StrDict errMsgs[];
	// External handle functions
	friend cl_int  initOpenCL(host_info_id host);
	friend cl_int loadMesh(void *Ptr1,void *Ptr2);
	friend void    shutdownOpenCL(void);
private:
	// Internal containers for platforms and devices
	static std::vector<oclPlatform> _platforms;
	static std::vector<oclDevice>   _devices;
	// Self pointer
	static ocl * _selfp;
	static myLock _lock;

	//
	static void   lock() { _lock.lock(); }
	static void unlock() { _lock.unlock(); }

	static int invokeTestKernel(void);
  public:

	static ocl*& instance() { return _selfp; }
	// Dtr
	~ocl();

	// Returns the number of supported devices for current platform
	cl_uint isDeviceSupported(cl_device_type devType);
	cl_int loadProgram(const std::string& srcPath,bool quiet = true);

 protected:
	unsigned int _currPlatform; // Index of platform in use
	unsigned int _currDevice;   // index of device in use
	oclContext	_context;
	oclCmdQ		_queue;
	oclProgram	_program;
	oclKernel	_kernel;
	oclEvent 	_event;
	cl_int	  	_err;

 private:
	// Ctr
	explicit ocl(cl_uint devType = CL_DEVICE_TYPE_DEFAULT);

	// Not implemented
	ocl(const ocl&);
	ocl& operator=(const ocl&);
};


StrDict ocl::errMsgs[] = {
		{TEST_FAILED, "Test failed at startup!\n"}
};

std::vector<cl::Platform> ocl::_platforms;
std::vector<cl::Device>   ocl::_devices;
ocl * ocl::_selfp = NULL;
myLock ocl::_lock;



/*---------------------------------------------------------------------
 * Invoke kernel dummy for start working
 *--------------------------------------------------------------------*/
int ocl::invokeTestKernel(void)
{
	static const char srckernel[] =
	"__kernel void                       \n"
	"VecAdd(const __global float * M,    \n"
	"       const __global float * V,      \n"
	"             __global float * U,      \n"
    "        const unsigned int     N)      \n"
	"{                                   \n"
	"    size_t id =  get_global_id(0);  \n"
	"    if (id < N) {                   \n"
	"      U[id] = M[id] * V[id]; }         \n"
	"}                                   \n"
	;
	const cl_uint width = 10000U;
	const cl_uint mem_size = width * sizeof(cl_float);
	const cl_float3 values = { 0.5f, 2.0f, 1.0f };
	float M[width], V[width], U[width];
	fvFillArray(M, (cl_int)width, &values.s[0]);
	fvFillArray(V, (cl_int)width, &values.s[1]);
	// Get a list of supported platforms
	cl::Platform::get(& ocl::_platforms);

	// Pick first platform
	cl_context_properties cprops[] = {
			CL_CONTEXT_PLATFORM, (cl_context_properties)(_platforms[0])(), 0};
	cl::Context context(CL_DEVICE_TYPE_GPU, cprops);

	// Query the set of devices attched to the context
	std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();

	// Create and program from source
	cl::Program::Sources sources(1, std::make_pair(srckernel, 0));
	cl::Program program(context, sources);
	mfp_debug("Before building CL program\n");

	try {
		cl_int err = program.build(devices, "-Werror -cl-fast-relaxed-math");
	} catch (cl::Error err )
	{
		std::string log( program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(devices[0]));
		throw cl::Error(err.err(), log.c_str());
	}

	// Alloc memory and fill in arrays


	// Create buffer for M and copy host contents
	cl::Buffer MBuffer = cl::Buffer(
	            context,
	            CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
	            mem_size,
	            (void *) &M[0]);
	// Create buffer for V and copy host contents
	cl::Buffer VBuffer = cl::Buffer(
	            context,
	            CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
	            mem_size,
	            (void *) &V[0]);
	// Create buffer for U vector result
	cl::Buffer UBuffer = cl::Buffer(
	            context,
	            CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
	            mem_size,
	            (void *) &U[0]);

	// Create kernel object
	cl::Kernel kernel(program, "VecAdd");

	// Set kernel arguments
	kernel.setArg(0, MBuffer);
	kernel.setArg(1, VBuffer);
	kernel.setArg(2, UBuffer);
    kernel.setArg(3, width);

	// Create command queue
	cl::CommandQueue queue(context, devices[0], 0);

	// Do the work
	queue.enqueueNDRangeKernel(
	            kernel,
	            cl::NullRange,
	            cl::NDRange(width),
	            cl::NullRange);

	// Map UBuffer to host pointer. This enforces a sync with
	// the host backing space, remember we choose GPU device.
	float * output = (float *) queue.enqueueMapBuffer(
	            UBuffer,
	            CL_TRUE, // block
	            CL_MAP_READ,
	            0,
	            mem_size);

	// validate results
	for (cl_uint i = 0; i < width; ++i) {
		if (values.s[2] != U[i])
			throw cl::Error(TEST_FAILED, getString(errMsgs,FV_SIZEOF_ARRAY(errMsgs),TEST_FAILED));
	}

	// Finally release our hold on accessing the memory
	cl_int err = queue.enqueueUnmapMemObject(
	            UBuffer,
	            (void *) output);
	return(1);
}

ocl::ocl(cl_uint devType)
{
	// Check for OpenCL support
	if (ocl::_platforms.empty() == 0) {
		_err = oclPlatform::get(&_platforms);
	}

	// Find first platform which support
	_currPlatform = 0;
	// For the first platform get specific devices
	_err = _platforms.at(_currPlatform).getDevices(devType, &ocl::_devices);

	cl_context_properties properties[] =
	        { CL_CONTEXT_PLATFORM, (cl_context_properties)(_platforms[_currPlatform])(), 0};

	// Setup a context and a command queue
	_currDevice = 0;
	_context = oclContext(devType,properties);
	_queue = oclCmdQ(_context, ocl::_devices.at(_currDevice), 0, &_err);

}

ocl::~ocl()
{

}

cl_uint ocl::isDeviceSupported(cl_device_type devType)
{
	cl_uint numDevs = 0;
	for (size_t i = 0; i < _devices.size(); ++i) {
		if (devType == _devices[i].getInfo<CL_DEVICE_TYPE>()) ++numDevs;
	}
	return numDevs;
}


cl_int ocl::loadProgram(const std::string& srcPath,bool quiet)
{
	// Load OpenCL source
	std::ifstream srcfile(srcPath);
	std::string src(std::istreambuf_iterator<char>(srcfile),
			std::istreambuf_iterator<char>(0));
	try {
		// Compile the device program
		_program = oclProgram(_context, oclProgram::Sources(1,
			std::make_pair(src.c_str(), src.size())));
		_err = _program.build(_devices, "-Werror");
	} catch(oclError err) {
		_err = err.err();
		std::cerr  << "ERROR: "
				<< err.what()
				<< "("
				<< err.err()
				<< ")"
				<< std::endl;
	}

	if (! quiet) {
		std::cout << "Build Status: " << _program.getBuildInfo<CL_PROGRAM_BUILD_STATUS>(_devices[_currDevice]) << std::endl;
		std::cout << "Build Options:\t" << _program.getBuildInfo<CL_PROGRAM_BUILD_OPTIONS>(_devices[_currDevice]) << std::endl;
		std::cout << "Build Log:\t " << _program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(_devices[_currDevice]) << std::endl;
	}

	return _err;

}
//} // end namespace

cl_int initOpenCL(host_info_id host)
{
	/*assert((devType & CL_DEVICE_TYPE_GPU) ||
		   (devType & CL_DEVICE_TYPE_ACCELERATOR) ||
		   (devType & CL_DEVICE_TYPE_CPU) ||
		   (devType & CL_DEVICE_TYPE_CUSTOM));*/
	ocl::lock();

	cl_int err;

	if (! ocl::_selfp) {
		atexit(shutdownOpenCL);

		// Get all platforms
		cl::Platform::get(&ocl::_platforms);
		if (ocl::_platforms.empty()) return (-1);
	}



	// Get a name of first supported platform
	host->platform = ocl::_platforms.at(0).getInfo<CL_PLATFORM_NAME>();

	// Scan for available devices
	cl_int totalNumOfdevs = 0;
	try {
		// Single instance
		if (! ocl::_selfp) {
			ocl::_selfp = new ocl(CL_DEVICE_TYPE_ALL);
		}

		// Collect info about CL devices
		std::vector<device_info_t> * ptr = &host->devices;
		device_info_t dev;
		for (int i = 0; i < FV_SIZEOF_ARRAY(_devices); ++i) {
			dev.type = _devices[i];
			dev.num = ocl::_selfp->isDeviceSupported(dev.type);
			if (dev.num > 0) {
				ptr->push_back(dev);
				totalNumOfdevs += dev.num;
			}
		}
		// Inovke test
		err = ocl::invokeTestKernel();
	} catch (ocl::oclError& err) {
		std::cerr  << "ERROR: "
				   << err.what()
				   << "("
				   << err.err()
				   << ")"
				   << std::endl;
		err = -1;
	}
	ocl::unlock();
	return err;
}

cl_int loadMesh(void *Ptr1,void *Ptr2)
{
	static const CoordType cl_max = -std::numeric_limits<CoordType>::max();
	// Handle to singleton device
	ocl*& OCL(ocl::instance());
	assert(OCL != NULL);

	cl_float MinMax[2];
	// Handle to mesh
	FemViewer::Mesh *mesh_ = reinterpret_cast<FemViewer::Mesh *>(Ptr1);
	FemViewer::Field *field_ = reinterpret_cast<FemViewer::Field *>(Ptr2);
	// Get size of indices
	cl_uint totalSize = FemViewer::Mesh::SetElements(*mesh_);

	// Fill in with indices
	std::vector<cl_int> indices(totalSize,0);
	FemViewer::Mesh::FillArrayWithIndicesOfElemNodes(mesh_,indices.data());
	assert(indices.size() == totalSize);
	const cl_uint mem_IndSize = sizeof(cl_int) * totalSize;

	// Fill in coordinates
	std::vector<CoordType> vertices(3*mesh_->GetMaxNodeId(),cl_max);
	FemViewer::Mesh::FillArrayWithElemCoords(mesh_,vertices.data());
	const cl_uint mem_VertSize = sizeof(cl_float) * vertices.size();

	// Fill in solution coefficient
	cl_uint nCoeffs= field_->CountCoeffs();
	assert(nCoeffs > 0 && nCoeffs <= mesh_->GetElems().size()*450);
	std::vector<cl_float> coeffs(nCoeffs,-10e200);
	cl_uint coeffSize = FemViewer::Field::FillArrayWithCoeffs(field_,coeffs.data());
	assert(coeffSize == coeffs.size());
	const cl_uint mem_CoeffSize = sizeof(cl_float) * coeffSize;

	cl_context_properties cprops[] = {
				CL_CONTEXT_PLATFORM, (cl_context_properties)(OCL->_platforms[0])(), 0};
	cl::Context context(CL_DEVICE_TYPE_GPU, cprops);

	cl::Buffer IndicesBuffer = cl::Buffer(
				context,
				CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
				mem_IndSize,
				(void *) indices.data());

	cl::Buffer CoordBuffer = cl::Buffer(
			context,
			CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
			mem_VertSize,
			(void *) vertices.data());

	cl::Buffer CoeffBuffer = cl::Buffer(
			context,
			CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
			mem_CoeffSize,
			(void *) coeffs.data());

	cl::Buffer MnMxBuffer = cl::Buffer(
			context,
			CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
		    2 * sizeof(cl_float),
		    (void *) &MinMax[0]);

	// Load OpenCL kernel
	std::ifstream srcfile(kernels[1]);
	std::string src(std::istreambuf_iterator<char>(srcfile),
						std::istreambuf_iterator<char>(0));

	// Compile the device program
	cl::Program prog(context, cl::Program::Sources(1,
				std::make_pair(src.c_str(), src.size())));
	prog.build(ocl::_devices, "-Werror");

	cl::Kernel kernel = cl::Kernel(prog, "MinMaxRenge");
	kernel.setArg(0, CoordBuffer);
	kernel.setArg(1, IndicesBuffer);
	kernel.setArg(2, CoeffBuffer);
	kernel.setArg(3, MnMxBuffer);
	kernel.setArg(4, mesh_->GetElems().size());

	// Create command queue
	cl::CommandQueue queue(context, OCL->_devices[0], 0);
	queue.enqueueWriteBuffer(IndicesBuffer,CL_TRUE,0,mem_IndSize, indices.data());
	queue.enqueueWriteBuffer(CoordBuffer,CL_TRUE,0,mem_VertSize, vertices.data());
	queue.enqueueWriteBuffer(CoeffBuffer,CL_TRUE,0,mem_CoeffSize, coeffs.data());

	// Do the work
	queue.enqueueNDRangeKernel(
		            kernel,
		            cl::NullRange,
		            cl::NDRange(mesh_->GetElems().size()),
		            cl::NullRange);

	queue.finish();
	queue.enqueueReadBuffer(MnMxBuffer,CL_TRUE,0,2*sizeof(cl_float), &MinMax[0]);

	return 0;
}


void shutdownOpenCL(void)
{
	mfp_debug("OpenCL: shutdown\n");
	ocl::lock();
	if (ocl::_selfp != 0) { delete ocl::_selfp; ocl::_selfp = 0; }
	ocl::unlock();
}







